
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <windows.h>
#include <stdio.h>

// kernels
#include "common.cuh"
#include "sgemm_kernel.cuh"




// Matrix dimensions
// (chosen as multiples of the thread block size for simplicity)
//#define WA (20 * BLOCK_SIZE_X) // Matrix A width
//#define HA (20 * BLOCK_SIZE_Y) // Matrix A height
//#define WB (20 * BLOCK_SIZE_Y) // Matrix B width

#define N 64
#define WA N
#define HA N
#define WB N

#define HB WA  // Matrix B height
#define WC WB  // Matrix C width
#define HC HA  // Matrix C height

////////////////////////////////////////////////////////////////////////////////
// declaration, forward
int runTest(int argc, char** argv);
bool shouldGetUserInput(int argc, char** argv);
void randomInit(_float*, int);
void printDiff(float*, float*, int, int);

void computeGold(float*, const float*, const float*, unsigned int, unsigned int, unsigned int);
int PrintDevices(int deviceCount, int deviceSelected);

void transpose(_float* A, int width, int height)
{
	 _float* tmp = (_float*) malloc(width*height*sizeof(_float));
	 for(int i=0; i < height; ++i)
		 for(int j=0; j < width; ++j)
			 tmp[j*height + i] = A[i*width + j];

	for(int i=0; i < width*height; ++i)		
		A[i] = tmp[i];

	free(tmp);
}

int PrintDevices(int deviceCount, int deviceSelected)
{
    cudaError_t err = cudaSuccess;

    cudaDeviceProp deviceProperty;
    for (int currentDeviceId = 0; currentDeviceId < deviceCount; ++currentDeviceId)
    {
        memset(&deviceProperty, 0, sizeof(cudaDeviceProp));
        err = cudaGetDeviceProperties(&deviceProperty, currentDeviceId);
        CheckConditionXR_(err == cudaSuccess, err);

        printf("\ndevice name: %s", deviceProperty.name);
        if (currentDeviceId == deviceSelected)
        {
            printf("    <----- creating CUcontext on this");    
        }
        printf("\n");

        printf("device sharedMemPerBlock: %Iu \n", deviceProperty.sharedMemPerBlock);
        printf("device totalGlobalMem: %Iu \n", deviceProperty.totalGlobalMem);
        printf("device regsPerBlock: %d \n", deviceProperty.regsPerBlock);
        printf("device warpSize: %d \n", deviceProperty.warpSize);
        printf("device memPitch: %Iu \n", deviceProperty.memPitch);
        printf("device maxThreadsPerBlock: %d \n", deviceProperty.maxThreadsPerBlock);
        printf("device maxThreadsDim[0]: %d \n", deviceProperty.maxThreadsDim[0]);
        printf("device maxThreadsDim[1]: %d \n", deviceProperty.maxThreadsDim[1]);
        printf("device maxThreadsDim[2]: %d \n", deviceProperty.maxThreadsDim[2]);
        printf("device maxGridSize[0]: %d \n", deviceProperty.maxGridSize[0]);
        printf("device maxGridSize[1]: %d \n", deviceProperty.maxGridSize[1]);
        printf("device maxGridSize[2]: %d \n", deviceProperty.maxGridSize[2]);
        printf("device totalConstMem: %Iu \n", deviceProperty.totalConstMem);
        printf("device major: %d \n", deviceProperty.major);
        printf("device minor: %d \n", deviceProperty.minor);
        printf("device clockRate: %d \n", deviceProperty.clockRate);
        printf("device textureAlignment: %Iu \n", deviceProperty.textureAlignment);
        printf("device deviceOverlap: %d \n", deviceProperty.deviceOverlap);
        printf("device multiProcessorCount: %d \n", deviceProperty.multiProcessorCount);

        printf("\n");
    }

    return cudaSuccess;
}

int main(int argc, char** argv)
{
	int err = runTest(argc, argv);
    CheckConditionXR(err == cudaSuccess, -1);

    bool getUserInput = shouldGetUserInput(argc, argv);

    if (getUserInput)
    {
        printf("Hit any key to terminate\n");
        getchar();
    }

    return 0;
}


////////////////////////////////////////////////////////////////////////////////
//! Run a simple test for CUDA
////////////////////////////////////////////////////////////////////////////////
int
runTest(int argc, char** argv)
{
    cudaError_t err = cudaSuccess;

    int deviceCount = 0;
    err = cudaGetDeviceCount(&deviceCount);
    CheckConditionXR_(err == cudaSuccess, err)

    if (deviceCount == 0) {
        printf("error: no devices supporting CUDA.\n");
        return(-1);
    }

    int deviceOrdinal = 0;
    if (argc > 1)
    {
        deviceOrdinal = atoi(argv[1]);
    }

    CheckConditionXR_(deviceOrdinal <= deviceCount, cudaErrorInvalidDevice);

    err = (cudaError_t)PrintDevices(deviceCount, deviceOrdinal);
    CheckConditionXR_(err == cudaSuccess, err)

    err = cudaSetDevice(deviceOrdinal);
    CheckConditionXR_(err == cudaSuccess, err)

    // set seed for rand()
    srand(2006 );

    // allocate host memory for matrices A and B
    unsigned int size_A = WA * HA;
    unsigned int mem_size_A = sizeof(_float) * size_A;
    _float* h_A = (_float*) malloc(mem_size_A);
    unsigned int size_B = WB * HB;
    unsigned int mem_size_B = sizeof(_float) * size_B;
    _float* h_B = (_float*) malloc(mem_size_B);

    // initialize host memory
    randomInit(h_A, size_A);
    randomInit(h_B, size_B);

	transpose(h_A, WA, HA);
	//	for(int i=0; i < HC; ++i)
	//{
	//	for(int j=0; j < WC; ++j)
	//		printf("(%i,%i) %f, ", i, j, h_A[i*WC+j]);
	//	printf("\n");
	//}
	transpose(h_B, WB, HB);

    // allocate device memory
    _float* A;
    cudaMalloc((void**) &A, mem_size_A);
    _float* B;
    cudaMalloc((void**) &B, mem_size_B);

    // copy host memory to device
    cudaMemcpy(A, h_A, mem_size_A, cudaMemcpyHostToDevice);
    cudaMemcpy(B, h_B, mem_size_B, cudaMemcpyHostToDevice);

    // allocate device memory for result
    unsigned int size_C = WC * HC;
    unsigned int mem_size_C = sizeof(float) * size_C;
    float* C;
    cudaMalloc((void**) &C, mem_size_C);

    // allocate host memory for the result
    float* h_C = (float*) malloc(mem_size_C);
    
    // create and start timer
  

    // block size and grid size
    dim3 blockSize(BLOCK_X, BLOCK_Y);
    //dim3 grid(WC / blockSize.x, HC / blockSize.y);
	dim3 grid(HC / VEC_LENGTH, WC / BY);
	printf("grid.x : %i , grid.y %i\n", grid.x, grid.y);
	// start time
	cudaEvent_t start, stop; 
	float processTimeMS; 
	cudaEventCreate(&start); 
	cudaEventCreate(&stop); 
	cudaEventRecord( start, 0 );
    // execute the kernel
    //mulMatrix<<< grid, threads >>>(A, B, C, WA, WB);
	SGEMM<<< grid, blockSize >>>(A, B, C, HA, WA);	
	cudaDeviceSynchronize();
	//finish time	
	cudaEventRecord( stop, 0 ); 
	cudaEventSynchronize( stop ); 
	cudaEventElapsedTime( &processTimeMS, start, stop ); 
	cudaEventDestroy( start ); 
	cudaEventDestroy( stop );

    CheckConditionXR_(cudaGetLastError() == cudaSuccess, -1);


    // copy result from device to host
    err = cudaMemcpy(h_C, C, mem_size_C, cudaMemcpyDeviceToHost);
    CheckConditionXR_(err == cudaSuccess, err)


    printf("Processing time: %f (ms) \n", processTimeMS);
    double gflops = ((2 * WA)/1000.0) * (WC/1000.0) * (HC/1000.0);
	double gFlopsPerS = gflops / (processTimeMS/(double)1000.0);
	printf("Width(A) = %i, Width(C) = %i, Height(C) = %i\n", WA, WC, HC);
	printf("Glops/s: %f\n", gFlopsPerS);

	//transpose C back to row-major
	//transpose(h_C, HC, WC);

    // compute reference solution
    float* reference = (float*) malloc(mem_size_C);
    computeGold(reference, h_A, h_B, HA, WA, WB);

    float acc = 0;
	bool res = true;
    for (unsigned int i = 0; i < size_C; ++i)
    {
        float temp = reference[i] - h_C[i];
		if(temp >= 1e-4f)
			res = false;
        float temp_sq = temp * temp;
        acc += temp_sq;
    }
   // bool res = (acc <= 1e-4f);

    // check result
    printf("Test %s , acc = %f \n", (1 == res) ? "PASSED" : "FAILED", acc);
	//
	for(int i=0; i < HC; ++i)
	{
		for(int j=0; j < WC; ++j)
			printf("(%i,%i) %f, %f", i, j, reference[i*WC+j], h_C[i*WC+j]);
		printf("\n");
	}
    //if (res!=1) printDiff(reference, h_C, WC, HC);

    // clean up memory
    free(h_A);
    free(h_B);
    free(h_C);
    free(reference);
    cudaFree(A);
    cudaFree(B);
    cudaFree(C);

    cudaThreadExit();

    return cudaSuccess;
}

// Allocates a matrix with random float entries.
void randomInit(_float* data, int size)
{
    for (int i = 0; i < size; ++i)
       data[i] = rand() / (float)RAND_MAX;
	   //data[i] = 1.0;

	//for(int i=0; i < HC; ++i)
	//{
	//	for(int j=0; j < WC; ++j)
	//		data[i*WC+j] = j;
	//	
	//}
}

void printDiff(float *data1, float *data2, int width, int height)
{
  int i,j,k;
  int error_count=0;
  for (j=0; j<height; j++) {
    for (i=0; i<width; i++) {
      k = j*width+i;
      if (data1[k] != data2[k]) {
         printf("diff(%d,%d) CPU=%4.4f, GPU=%4.4f n", i,j, data1[k], data2[k]);
         error_count++;
      }
    }
  }
  printf(" nTotal Errors = %d n", error_count);
}

bool shouldGetUserInput(int argc, char** argv)
{
    for(int i = 1; i < argc; i++)
    {
        if (strcmp(argv[i], "noprompt") == 0)
        {
            return false;
        }
    }

    return true;
}